#include <cutil_inline.h>
#include "sortingalgorithms.h"

/*
 * Tady smer resit nemusime, vzdy je rostouci.
 */
__device__ inline void compareAndExchange(int &x, int &y)
{
    if (x > y)
    {
	int tmp = y;
	y = x;
	x = tmp;
    }
}

/*
 * Serazeni pole, ktere se vejde do shared pameti.
 */
__global__ void evenOddMergeSortImpl(int * items, unsigned int cnt, unsigned long offset = 0)
{
    // Prekopirovani odpovidajicich casti posloupnosti do shared pameti
    __shared__ int cudaItems[SHARED_SIZE];
    items += blockIdx.x * SHARED_SIZE + threadIdx.x + offset;
    cudaItems[threadIdx.x] = items[0];
    cudaItems[threadIdx.x + (SHARED_SIZE / 2)] = items[SHARED_SIZE / 2];

    // Vnejsi cyklus - pocet radicich fazi
    for (unsigned int i = 2; i <= cnt; i <<= 1)
    {
	unsigned int j = i / 2;
	unsigned int offset = threadIdx.x % j;

	__syncthreads();
	unsigned int position = 2 * threadIdx.x - (threadIdx.x % j);
	compareAndExchange(cudaItems[position], cudaItems[position + j]);
	j >>= 1;
	// Rekurzivni merge
	for (; j > 0; j >>= 1)
	{
	    __syncthreads();
	    position = 2 * threadIdx.x - (threadIdx.x % j);
	    if (offset >= j)
	    {
		compareAndExchange(cudaItems[position - j], cudaItems[position]);
	    }
	}
    }
    __syncthreads();
    items[0] = cudaItems[threadIdx.x];
    items[SHARED_SIZE / 2] = cudaItems[threadIdx.x + (SHARED_SIZE / 2)];
}

/*
 * Serazeni pole, kdyz se nevejde do shared pameti, musi se pracovat primo s puvodnim (pomalejsi).
 * items - Pole prvku k razeni
 * cnt - pocet prvku pole
 * size - velikost prave razene podposloupnosti
 * j - velikost prave mergovane podposloupnosti
 */
__global__ void evenOddMergeSortImplBig(int * items, unsigned long cnt, unsigned long size, unsigned int j, unsigned long offset)
{
    unsigned int threadPos = blockIdx.x * blockDim.x + threadIdx.x + offset;
    unsigned int position = 2 * threadPos - (threadPos % j);
    if (j < size / 2)
    {
	unsigned int off = threadPos & ( (size / 2) - 1);
	if (off >= j)
	{
	    compareAndExchange(items[position - j], items[position]);
	}
    } else
    {
	compareAndExchange(items[position], items[position + j]);
    }
}

// Sort items with EvenOddMergeSort algorithm
extern "C" void evenOddMergeSort(int * items, unsigned long cnt, unsigned int blockCnt){
    if (cnt < 2)
	return;
    unsigned int threadCnt = SHARED_SIZE / 2;
    if (cnt <= SHARED_SIZE)
    {
	evenOddMergeSortImpl<<<blockCnt, threadCnt>>>(items, cnt);
    } else
    {
	unsigned int l = blockCnt * threadCnt * 2;
	// Nejprve jeden merge casti pole, ktera se jeste vejde do pameti
        for (unsigned long x = 0; x <= cnt; x += (l + SHARED_SIZE / 2))
	{
	    evenOddMergeSortImpl<<<blockCnt, threadCnt>>>(items, SHARED_SIZE, x);
	}
	// Odted se musi mergovat uz primo nad puvodnim polem
	for (unsigned long i = 2 * SHARED_SIZE; i <= cnt; i <<= 1)
	{
	    // Rekurzivni merge rozepsany na for cyklus
	    for (unsigned long j = i / 2; j > 0; j >>= 1)
	    {
		for (unsigned long x = 0; x <= cnt; x += (l + j))
		{
		    evenOddMergeSortImplBig<<<blockCnt, threadCnt>>>(items, cnt, i, j, x);
		}
	    }
	}
    }
}
